ebab7d0da53c51438720d60c6e2e467c - /page/wikibook/entry.mu

Micron Document

Livres et Wikis | Archives | Info

🔍Search

CUDA

99.5 KB · ⇊ 26 parts · ⤓ Micron 99.5 KB · ⤓ HTML 321.4 KB

layout: Wide · Narrow · Centered

──────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────

top

CUDA is a proprietarycite-ref-0-2-0[2] parallel computing platform and application programming interface (API) that allows software to use certain types of graphics processing units (GPUs) for accelerated general-purpose processing, significantly broadening their utility in scientific and high-performance computing. CUDA was created by Nvidia starting in 2004 and was officially released by in 2007.cite-ref-3[3] When it was first introduced, the name was an acronym for Compute Unified Device Architecture,cite-ref-cuda-intro-anandtech-4-0[4] but Nvidia later dropped the common use of the acronym and now rarely expands it.cite-ref-5[5]

CUDA is both a software layer that manages data, giving direct access to the GPU and CPU as necessary and a library of APIs that enable parallel computation for various needs.cite-ref-cuda-intro-tomshardware-6-0[6]cite-ref-7[7] In addition to drivers and runtime kernels, the CUDA platform includes compilers, libraries and developer tools to help programmers accelerate their applications.

CUDA is written in C but is designed to work with a wide array of other programming languages including C++, Fortran, Python and Julia. This accessibility makes it easier for specialists in parallel programming to use GPU resources, in contrast to prior APIs like Direct3D and OpenGL, which require advanced skills in graphics programming.cite-ref-8[8] CUDA-powered GPUs also support programming frameworks such as OpenMP, OpenACC and OpenCL.cite-ref-9[9]cite-ref-cuda-intro-tomshardware-6-1[6]

Contents

• Background

• Ontology

• Programming abilities

• Advantages

• Limitations

• Example

• GPUs supported

• Version features and specifications

• Data types

• Tensor cores

• Technical specifications

• Multiprocessor architecture

• Usages of CUDA architecture

• Comparison with competitors

• Intel OneAPI

• AMD ROCm

• See also

• References

• Further reading

• External links

Background

The graphics processing unit (GPU), as a specialized computer processor, addresses the demands of real-time high-resolution 3D graphics compute-intensive tasks. By 2012, GPUs had evolved into highly parallel multi-core systems allowing efficient manipulation of large blocks of data. This design is more effective than general-purpose central processing unit (CPUs) for algorithms in situations where processing large blocks of data is done in parallel, such as:

• cryptographic hash functions

• machine learning

• molecular dynamics simulations

• physics engines

The origins of CUDA trace back to the early 2000s, when Ian Buck, a computer science Ph.D. student at Stanford University, began experimenting with using GPUs for purposes beyond rendering graphics. Buck had first become interested in GPUs during his undergraduate studies at Princeton University, initially through video gaming. After graduation, he interned at Nvidia, gaining deeper exposure to GPU architecture. At Stanford, he built an 8K gaming rig using 32 GeForce graphics cards, originally to push the limits of graphics performance in games like Quake and Doom. However, his interests shifted toward exploring the potential of GPUs for general-purpose parallel computing.cite-ref-1-10-0[10]

To that end, Buck developed Brook, a programming language designed to enable general-purpose computing on GPUs. His work attracted support from both Nvidia and the Defense Advanced Research Projects Agency (DARPA). In 2004, Nvidia hired Buck and paired him with John Nickolls, the company’s director of architecture for GPU computing. Together, they began transforming Brook into what would become CUDA.cite-ref-1-10-1[10] CUDA was officially released by Nvidia in 2007.

Under the leadership of Nvidia CEO Jensen Huang, CUDA became central to the company’s strategy of positioning GPUs as versatile hardware for scientific applications. By 2015, CUDA’s development increasingly focused on accelerating machine learning and artificial neural network workloads.cite-ref-11[11]

Ontology

The following table offers a non-exact description for the ontology of the CUDA framework.

| memory (hardware) | memory (code, or variable scoping ) |

|---|---|

| RAM | non-CUDA variables |

| VRAM , GPU L2 cache | global, const, texture |

| GPU L1 cache | local, shared |

| | |

| GPU L0 cache, register | |

| memory (hardware) | computation (hardware) | computation (code syntax) |

|---|---|---|

| RAM | host | program |

| VRAM , GPU L2 cache | device | grid |

| GPU L1 cache | SM ("streaming multiprocessor") | block |

| | warp = 32 threads | |

| GPU L0 cache, register | thread (aka. "SP", "streaming processor… | |

| memory (hardware) | computation (code semantics) |

|---|---|

| RAM | one routine call |

| VRAM , GPU L2 cache | simultaneous call of the same subroutin… |

| GPU L1 cache | individual subroutine call |

| | SIMD instructions |

| GPU L0 cache, register | analogous to individual scalar ops with… |

Programming abilities

The CUDA platform is accessible to software developers through CUDA-accelerated libraries, compiler directives such as OpenACC, and extensions to industry-standard programming languages including C, C++, Fortran and Python. C/C++ programmers can use 'CUDA C/C++', compiled to PTX with nvcc, Nvidia's LLVM-based C/C++ compiler, or by clang itself.cite-ref-12[12] Fortran programmers can use 'CUDA Fortran', compiled with the PGI CUDA Fortran compiler from The Portland Group. Python programmers can use the cuNumeric library to accelerate applications on Nvidia GPUs.

In addition to libraries, compiler directives, CUDA C/C++ and CUDA Fortran, the CUDA platform supports other computational interfaces, including the Khronos Group's OpenCL,cite-ref-13[13] Microsoft's DirectCompute, OpenGL Compute Shader and C++ AMP.cite-ref-14[14] Third party wrappers are also available for Python, Perl, Fortran, Java, Ruby, Lua, Common Lisp, Haskell, R, MATLAB, IDL, Julia, and native support in Mathematica.

In the computer game industry, GPUs are used for graphics rendering, and for game physics calculations (physical effects such as debris, smoke, fire, fluids); examples include PhysX and Bullet. CUDA has also been used to accelerate non-graphical applications in computational biology, cryptography and other fields by an order of magnitude or more.cite-ref-ioannidis08-15-0[15]cite-ref-16[16]cite-ref-manavski2008-17-0[17]cite-ref-18[18]cite-ref-19[19]

CUDA provides both a low level API (CUDA Driver API, non single-source) and a higher level API (CUDA Runtime API, single-source). The initial CUDA SDK was made public on 15 February 2007, for Microsoft Windows and Linux. Mac OS X support was later added in version 2.0,cite-ref-20[20] which supersedes the beta released February 14, 2008.cite-ref-21[21] CUDA works with all Nvidia GPUs from the G8x series onwards, including GeForce, Quadro and the Tesla line. CUDA is compatible with most standard operating systems.

CUDA 8.0 comes with the following libraries (for compilation & runtime, in alphabetical order):

• cuBLAS – CUDA Basic Linear Algebra Subroutines library

• CUDART – CUDA Runtime library

• cuFFT – CUDA Fast Fourier Transform library

• cuRAND – CUDA Random Number Generation library

• cuSOLVER – CUDA based collection of dense and sparse direct solvers

• cuSPARSE – CUDA Sparse Matrix library

• NPP – NVIDIA Performance Primitives library

• nvGRAPH – NVIDIA Graph Analytics library

• NVML – NVIDIA Management Library

• NVRTC – NVIDIA Runtime Compilation library for CUDA C++

CUDA 8.0 comes with these other software components:

• nView – NVIDIA nView Desktop Management Software

• NVWMI – NVIDIA Enterprise Management Toolkit

• GameWorks PhysX – is a multi-platform game physics engine

CUDA 9.0–9.2 comes with these other components:

• CUTLASS 1.0 – custom linear algebra algorithms,

• NVIDIA Video Decoder was deprecated in CUDA 9.2; it is now available in NVIDIA Video Codec SDK

CUDA 10 comes with these other components:

• nvJPEG – Hybrid (CPU and GPU) JPEG processing

CUDA 11.0–11.8 comes with these other components:cite-ref-22[22]cite-ref-23[23]cite-ref-24[24]cite-ref-25[25]

• CUB is new one of more supported C++ libraries

• MIG multi instance GPU support

• nvJPEG2000 – JPEG 2000 encoder and decoder

Advantages

CUDA has several advantages over traditional general-purpose computation on GPUs (GPGPU) using graphics APIs:

• Scattered reads – code can read from arbitrary addresses in memory.

• Unified virtual memory (CUDA 4.0 and above)

• Unified memory (CUDA 6.0 and above)

• Shared memory – CUDA exposes a fast shared memory region that can be shared among threads. This can be used as a user-managed cache, enabling higher bandwidth than is possible using texture lookups.cite-ref-26[26]

• Faster downloads and readbacks to and from the GPU

• Full support for integer and bitwise operations, including integer texture lookups

Limitations

• Whether for the host computer or the GPU device, all CUDA source code is now processed according to C++ syntax rules.cite-ref-cuda-prog-v8-27-0[27] This was not always the case. Earlier versions of CUDA were based on C syntax rules.cite-ref-28[28] As with the more general case of compiling C code with a C++ compiler, it is therefore possible that old C-style CUDA source code will either fail to compile or will not behave as originally intended.

• Interoperability with rendering languages such as OpenGL is one-way, with OpenGL having access to registered CUDA memory but CUDA not having access to OpenGL memory.

• Copying between host and device memory may incur a performance hit due to system bus bandwidth and latency (this can be partly alleviated with asynchronous memory transfers, handled by the GPU's DMA engine).

• Threads should be running in groups of at least 32 for best performance, with total number of threads numbering in the thousands. Branches in the program code do not affect performance significantly, provided that each of 32 threads takes the same execution path; the SIMD execution model becomes a significant limitation for any inherently divergent task (e.g. traversing a space partitioning data structure during ray tracing).

• No emulation or fallback functionality is available for modern revisions.

• Valid C++ may sometimes be flagged and prevent compilation due to the way the compiler approaches optimization for target GPU device limitations.

• C++ run-time type information (RTTI) and C++-style exception handling are only supported in host code, not in device code.

• In single-precision on first generation CUDA compute capability 1.x devices, denormal numbers are unsupported and are instead flushed to zero, and the precision of both the division and square root operations are slightly lower than IEEE 754-compliant single precision math. Devices that support compute capability 2.0 and above support denormal numbers, and the division and square root operations are IEEE 754 compliant by default. However, users can obtain the prior faster gaming-grade math of compute capability 1.x devices if desired by setting compiler flags to disable accurate divisions and accurate square roots, and enable flushing denormal numbers to zero.cite-ref-29[29]

• Unlike OpenCL, CUDA-enabled GPUs are only available from Nvidia as it is proprietary.cite-ref-cuda-products-30-0[30]cite-ref-0-2-1[2] Attempts to implement CUDA on other GPUs include:

• Project Coriander: Converts CUDA C++11 source to OpenCL 1.2 C. A fork of CUDA-on-CL intended to run TensorFlow.cite-ref-31[31]cite-ref-32[32]cite-ref-33[33]

• CU2CL: Convert CUDA 3.2 C++ to OpenCL C.cite-ref-34[34]

• GPUOpen HIP: A thin abstraction layer on top of CUDA and ROCm intended for AMD and Nvidia GPUs. Has a conversion tool for importing CUDA C++ source. Supports CUDA 4.0 plus C++11 and float16.

• ZLUDA is a drop-in replacement for CUDA on AMD GPUs and formerly Intel GPUs with near-native performance.cite-ref-35[35] The developer, Andrzej Janik, was separately contracted by both Intel and AMD to develop the software in 2021 and 2022, respectively. However, neither company decided to release it officially due to the lack of a business use case. AMD's contract included a clause that allowed Janik to release his code for AMD independently, allowing him to release the new version that only supports AMD GPUs.cite-ref-36[36]

• chipStar can compile and run CUDA/HIP programs on advanced OpenCL 3.0 or Level Zero platforms.cite-ref-37[37]

Example

This example code in C++ loads a texture from an image into an array on the GPU:

texture<float, 2, cudaReadModeElementType> tex;

void foo()

{

cudaArray* cu_array;

// Allocate array

cudaChannelFormatDesc description = cudaCreateChannelDesc<float>();

cudaMallocArray(&cu_array, &description, width, height);

// Copy image data to array

cudaMemcpyToArray(cu_array, image, width*height*sizeof(float), cudaMemcpyHostToDevice);

// Set texture parameters (default)

tex.addressMode[0] = cudaAddressModeClamp;

tex.addressMode[1] = cudaAddressModeClamp;

tex.filterMode = cudaFilterModePoint;

tex.normalized = false; // do not normalize coordinates

// Bind the array to the texture

cudaBindTextureToArray(tex, cu_array);

// Run kernel

dim3 blockDim(16, 16, 1);

dim3 gridDim((width + blockDim.x - 1)/ blockDim.x, (height + blockDim.y - 1) / blockDim.y, 1);

kernel<<< gridDim, blockDim, 0 >>>(d_data, height, width);

// Unbind the array from the texture

cudaUnbindTexture(tex);

} //end foo()

__global__ void kernel(float* odata, int height, int width)

{

unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;

unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

if (x < width && y < height) {

float c = tex2D(tex, x, y);

odata[y*width+x] = c;

}

Below is an example given in Python that computes the product of two arrays on the GPU. The unofficial Python language bindings can be obtained from PyCUDA.cite-ref-38[38]

import pycuda.compiler as comp

import pycuda.driver as drv

import numpy

import pycuda.autoinit

mod = comp.SourceModule(

"""

__global__ void multiply_them(float *dest, float *a, float *b)

{

const int i = threadIdx.x;

dest[i] = a[i] * b[i];

}

"""

)

multiply_them = mod.get_function("multiply_them")

a = numpy.random.randn(400).astype(numpy.float32)

b = numpy.random.randn(400).astype(numpy.float32)

dest = numpy.zeros_like(a)

multiply_them(drv.Out(dest), drv.In(a), drv.In(b), block=(400, 1, 1))

print(dest - a * b)

Additional Python bindings to simplify matrix multiplication operations can be found in the program pycublas.cite-ref-39[39]

import numpy

from pycublas import CUBLASMatrix

A = CUBLASMatrix(numpy.mat([[1, 2, 3], [4, 5, 6]], numpy.float32))

B = CUBLASMatrix(numpy.mat([[2, 3], [4, 5], [6, 7]], numpy.float32))

C = A * B

print(C.np_mat())

while CuPy directly replaces NumPy:cite-ref-40[40]

import cupy

a = cupy.random.randn(400)

b = cupy.random.randn(400)

dest = cupy.zeros_like(a)

print(dest - a * b)

GPUs supported

Supported CUDA compute capability versions for CUDA SDK version and microarchitecture (by code name):

|---|---|---|---|---|---|---|

| 1.0 | 1.0 – 1.1 | | | | | |

| 1.1 | 1.0 – 1.1+x | | | | | |

| 2.0 | 1.0 – 1.1+x | | | | | |

| 2.1 – 2.3.1 | 1.0 – 1.3 | | | | | |

| 3.0 – 3.1 | 1.0 | 2.0 | | | | |

| 3.2 | 1.0 | 2.1 | | | | |

| 4.0 – 4.2 | 1.0 | 2.1 | | | | |

| 5.0 – 5.5 | 1.0 | | 3.0 | 3.5 | | |

| 6.0 | 1.0 | | 3.2 | 3.5 | | |

| 6.5 | 1.1 | | | 3.7 | 5.x | |

| 7.0 – 7.5 | | 2.0 | | | 5.x | |

| 8.0 | | 2.0 | | | | 6.x |

| 9.0 – 9.2 | | | 3.0 | | | |

| 10.0 – 10.2 | | | 3.0 | | | |

| 11.0 | | | | 3.5 | | |

| 11.1 – 11.4 | | | | 3.5 | | |

| 11.5 – 11.7.1 | | | | 3.5 | | |

| 11.8 | | | | 3.5 | | |

| 12.0 – 12.6 | | | | | 5.0 | |

| 12.8 | | | | | 5.0 | |

| 12.9 | | | | | 5.0 | |

|---|---|---|---|---|---|---|

| 9.0 – 9.2 | 7.0 – 7.2 | | | | | |

| 10.0 – 10.2 | | 7.5 | | | | |

| 11.0 | | | 8.0 | | | |

| 11.1 – 11.4 | | | 8.6 | | | |

| 11.5 – 11.7.1 | | | 8.7 | | | |

| 11.8 | | | | 8.9 | 9.0 | |

| 12.0 – 12.6 | | | | | 9.0 | |

| 12.8 | | | | | | 12.0 |

| 12.9 | | | | | | 12.1 |

Note: CUDA SDK 10.2 is the last official release for macOS, as support will not be available for macOS in newer releases.

CUDA compute capability by version with associated GPU semiconductors and GPU card models (separated by their various application areas):

| Compute capability (version) | Micro- architecture | GPUs |

|---|---|---|

| 1.0 | Tesla | G80 |

| 1.1 | Tesla | G92, G94, G96, G98, G84, G86 |

| 1.2 | Tesla | GT218, GT216, GT215 |

| 1.3 | Tesla | GT200, GT200b |

| 2.0 | Fermi | GF100, GF110 |

| 2.1 | Fermi | GF104, GF106 GF108, GF114, GF116, GF117… |

| 3.0 | Kepler | GK104, GK106, GK107 |

| 3.2 | Kepler | GK20A |

| 3.5 | Kepler | GK110, GK208 |

| 3.7 | Kepler | GK210 |

| 5.0 | Maxwell | GM107, GM108 |

| 5.2 | Maxwell | GM200, GM204, GM206 |

| 5.3 | Maxwell | GM20B |

| 6.0 | Pascal | GP100 |

| 6.1 | Pascal | GP102, GP104, GP106, GP107, GP108 |

| 6.2 | Pascal | GP10B |

| 7.0 | Volta | GV100 |

| 7.2 | Volta | GV10B GV11B |

| 7.5 | Turing | TU102, TU104, TU106, TU116, TU117 |

| 8.0 | Ampere | GA100 |

| 8.6 | Ampere | GA102, GA103, GA104, GA106, GA107 |

| 8.7 | Ampere | GA10B |

| 8.9 | Ada Lovelace | AD102, AD103, AD104, AD106, AD107 |

| 9.0 | Hopper | GH100 |

| 10.0 | Blackwell | GB100 |

| 10.1 | Blackwell | |

| 10.3 | Blackwell | GB200, G10 |

| 12.0 | Blackwell | GB202, GB203, GB205, GB206, GB207 |

| 12.1 | Blackwell | |

| Compute capability (version) | Micro- architecture | GPUs |

| Compute capability (version) | GeForce |

|---|---|

| 1.0 | GeForce 8800 Ultra, GeForce 8800 GTX, G… |

| 1.1 | GeForce GTS 250, GeForce 9800 GX2, GeFo… |

| 1.2 | GeForce GT 340*, GeForce GT 330*, GeFor… |

| 1.3 | GeForce GTX 295, GTX 285, GTX 280, GeFo… |

| 2.0 | GeForce GTX 590, GeForce GTX 580, GeFor… |

| 2.1 | GeForce GTX 560 Ti, GeForce GTX 550 Ti,… |

| 3.0 | GeForce GTX 770, GeForce GTX 760, GeFor… |

| 3.5 | GeForce GTX Titan Z, GeForce GTX Titan… |

| 5.0 | GeForce GTX 750 Ti, GeForce GTX 750, Ge… |

| 5.2 | GeForce GTX Titan X, GeForce GTX 980 Ti… |

| 6.1 | Nvidia TITAN Xp, Titan X, GeForce GTX 1… |

| 7.0 | NVIDIA TITAN V |

| 7.5 | NVIDIA TITAN RTX, GeForce RTX 2080 Ti,… |

| 8.6 | GeForce RTX 3090 Ti, RTX 3090, RTX 3080… |

| 8.9 | GeForce RTX 4090, RTX 4080 Super, RTX 4… |

| 12.0 | GeForce RTX 5090, RTX 5080, RTX 5070 Ti… |

| Compute capability (version) | GeForce |

| Compute capability (version) | Quadro , NVS |

|---|---|

| 1.0 | Quadro FX 5600, Quadro FX 4600, Quadro… |

| 1.1 | Quadro FX 4700 X2, Quadro FX 3700, Quad… |

| 1.2 | Quadro FX 380 Low Profile, Quadro FX 18… |

| 1.3 | Quadro FX 5800, Quadro FX 4800, Quadro… |

| 2.0 | Quadro 6000, Quadro 5000, Quadro 4000,… |

| 2.1 | Quadro 2000, Quadro 2000D, Quadro 600,… |

| 3.0 | Quadro K5000, Quadro K4200, Quadro K400… |

| 3.5 | Quadro K6000, Quadro K5200 |

| 5.0 | Quadro K1200, Quadro K2200, Quadro K620… |

| 5.2 | Quadro M6000 24GB, Quadro M6000, Quadro… |

| 6.0 | Quadro GP100 |

| 6.1 | Quadro P6000, Quadro P5000, Quadro P400… |

| 7.0 | Quadro GV100 |

| 7.5 | Quadro RTX 8000, Quadro RTX 6000, Quadr… |

| 8.6 | RTX A6000, RTX A5500, RTX A5000, RTX A4… |

| 8.9 | RTX 6000 Ada, RTX 5880 Ada, RTX 5000 Ad… |

| 12.0 | RTX PRO 6000 Blackwell, RTX PRO 5000 Bl… |

| Compute capability (version) | Quadro , NVS |

| Compute capability (version) | Tesla/Datacenter |

|---|---|

| 1.0 | Tesla C870, Tesla D870, Tesla S870 |

| 1.3 | Tesla C1060, Tesla S1070, Tesla M1060 |

| 2.0 | Tesla C2075, Tesla C2050/C2070, Tesla M… |

| 3.0 | Tesla K10, GRID K340, GRID K520, GRID K2 |

| 3.5 | Tesla K40, Tesla K20x, Tesla K20 |

| 3.7 | Tesla K80 |

| 5.0 | Tesla M10 |

| 5.2 | Tesla M4, Tesla M40, Tesla M6, Tesla M60 |

| 6.0 | Tesla P100 |

| 6.1 | Tesla P40, Tesla P6, Tesla P4 |

| 7.0 | Tesla V100, Tesla V100S |

| 7.5 | Tesla T4 |

| 8.0 | A100 80GB, A100 40GB, A30 |

| 8.6 | A40, A16, A10, A2 |

| 8.9 | L40S, L40, L20, L4, L2 |

| 9.0 | H200, H100, GH200 |

| 10.0 | B200, B100, GB200 |

| 10.3 | B300, GB10 |

| 12.0 | B40 |

| Compute capability (version) | Tesla/Datacenter |

| Compute capability (version) | Tegra , Jetson , DRIVE |

|---|---|

| 3.2 | Tegra K1, Jetson TK1 |

| 5.3 | Tegra X1, Jetson TX1, Jetson Nano, DRIV… |

| 6.2 | Tegra X2, Jetson TX2, DRIVE PX 2 |

| 7.2 | Tegra Xavier, Jetson Xavier NX, Jetson… |

| 8.7 | Jetson Orin Nano, Jetson Orin NX, Jetso… |

| 10.1 | Jetson AGX Thor, DRIVE AGX Thor |

| Compute capability (version) | Tegra , Jetson , DRIVE |

* – OEM-only products

Version features and specifications

|---|---|---|---|---|

| Feature support (unlisted features are… | 1.0, 1.1 | 1.2, 1.3 | 2.x | 3.0 |

| Warp vote functions (__all(), __any()) | No | Yes | | |

| Warp vote functions (__ballot()) | No | | Yes | |

| Memory fence functions (__threadfence_s… | No | | Yes | |

| Synchronization functions (__syncthread… | No | | Yes | |

| Surface functions | No | | Yes | |

| 3D grid of thread blocks | No | | Yes | |

| Warp shuffle functions | No | | | Yes |

| Unified memory programming | No | | | Yes |

| Funnel shift | No | | | |

| Dynamic parallelism | No | | | |

| Uniform Datapath | No | | | |

| Hardware-accelerated async-copy | No | | | |

| Hardware-accelerated split arrive/wait… | No | | | |

| Warp-level support for reduction ops | No | | | |

| L2 cache residency management | No | | | |

| DPX instructions for accelerated dynami… | No | | | |

| Distributed shared memory | No | | | |

| Thread block cluster | No | | | |

| Tensor memory accelerator (TMA) unit | No | | | |

| Feature support (unlisted features are… | 1.0, 1.1 | 1.2, 1.3 | 2.x | 3.0 |

|---|---|---|---|---|

| Feature support (unlisted features are… | 3.2 | 3.5, 3.7, 5.x, 6.x, 7.0, 7.2 | 7.5 | 8.x |

| Funnel shift | Yes | | | |

| Dynamic parallelism | | Yes | | |

| Uniform Datapath | | | Yes | |

| Hardware-accelerated async-copy | | | | Yes |

| Hardware-accelerated split arrive/wait… | | | | Yes |

| Warp-level support for reduction ops | | | | Yes |

| L2 cache residency management | | | | Yes |

| Feature support (unlisted features are… | 3.2 | 3.5, 3.7, 5.x, 6.x, 7.0, 7.2 | 7.5 | 8.x |

| Feature support (unlisted features are… | |

|---|---|

| Feature support (unlisted features are… | 9.0, 10.x, 12.x |

| DPX instructions for accelerated dynami… | Yes |

| Distributed shared memory | Yes |

| Thread block cluster | Yes |

| Tensor memory accelerator (TMA) unit | Yes |

| Feature support (unlisted features are… | 9.0, 10.x, 12.x |

cite-ref-60[60]

Data types

Floating-point types

| Data type | Supported vector types | Storage Length Bits (complete vector) |

|---|---|---|

| E2M1 = FP4 | e2m1x2 / e2m1x4 | 8 / 16 |

| E2M3 = FP6 variant | e2m3x2 / e2m3x4 | 16 / 32 |

| E3M2 = FP6 variant | e3m2x2 / e3m2x4 | 16 / 32 |

| UE4M3 | ue4m3 | 8 |

| E4M3 = FP8 variant | e4m3 / e4m3x2 / e4m3x4 | 8 / 16 / 32 |

| E5M2 = FP8 variant | e5m2 / e5m2x2 / e5m2x4 | 8 / 16 / 32 |

| UE8M0 | ue8m0x2 | 16 |

| FP16 | f16 / f16x2 | 16 / 32 |

| BF16 | bf16 / bf16x2 | 16 / 32 |

| TF32 | tf32 | 32 |

| FP32 | f32 / f32x2 | 32 / 64 |

| FP64 | f64 | 64 |

|---|---|---|---|---|

| E2M1 = FP4 | 4 | 1 | 2 | 1 |

| E2M3 = FP6 variant | 6 | 1 | 2 | 3 |

| E3M2 = FP6 variant | 6 | 1 | 3 | 2 |

| UE4M3 | 7 | 0 | 4 | 3 |

| E4M3 = FP8 variant | 8 | 1 | 4 | 3 |

| E5M2 = FP8 variant | 8 | 1 | 5 | 2 |

| UE8M0 | 8 | 0 | 8 | 0 |

| FP16 | 16 | 1 | 5 | 10 |

| BF16 | 16 | 1 | 8 | 7 |

| TF32 | 19 | 1 | 8 | 10 |

| FP32 | 32 | 1 | 8 | 23 |

| FP64 | 64 | 1 | 11 | 52 |

| Data type | Comments |

|---|---|

| UE4M3 | Used for scaling (E2M1 only) |

| E5M2 = FP8 variant | Exponent/range of FP16, fits into 8 bits |

| UE8M0 | Used for scaling (any FP4 or FP6 or FP8… |

| BF16 | Exponent/range of FP32, fits into 16 bi… |

| TF32 | Exponent/range of FP32, mantissa/precis… |

Version support

| Data type | Basic Operations |

|---|---|

| 8-bit integer signed/unsigned | loading, storing, conversion |

| 16-bit integer signed/unsigned | general operations |

| 32-bit integer signed/unsigned | general operations |

| 64-bit integer signed/unsigned | general operations |

| any 128-bit trivially copyable type | general operations |

| 16-bit floating point FP16 | addition, subtraction, multiplication,… |

| 16-bit floating point BF16 | addition, subtraction, multiplication,… |

| 32-bit floating point | general operations |

| 32-bit floating point float2 and float4 | general operations |

| 64-bit floating point | general operations |

| Data type | Supported since | Atomic Operations |

|---|---|---|

| 8-bit integer signed/unsigned | 1.0 | — |

| 16-bit integer signed/unsigned | 1.0 | atomicCAS() |

| 32-bit integer signed/unsigned | 1.0 | atomic functions |

| 64-bit integer signed/unsigned | 1.0 | atomic functions |

| any 128-bit trivially copyable type | No | atomicExch, atomicCAS |

| 16-bit floating point FP16 | 5.3 | half2 atomic addition |

| 16-bit floating point FP16 | 5.3 | atomic addition |

| 16-bit floating point BF16 | 8.0 | atomic addition |

| 32-bit floating point | 1.0 | atomicExch() |

| 32-bit floating point | 1.0 | atomic addition |

| 32-bit floating point float2 and float4 | No | atomic addition |

| 64-bit floating point | 1.3 | atomic addition |

| Data type | Supported since for global memory |

|---|---|

| 8-bit integer signed/unsigned | — |

| 16-bit integer signed/unsigned | 3.5 |

| 32-bit integer signed/unsigned | 1.1 |

| 64-bit integer signed/unsigned | 1.2 |

| any 128-bit trivially copyable type | 9.0 |

| 16-bit floating point FP16 | 6.0 |

| 16-bit floating point FP16 | 7.0 |

| 16-bit floating point BF16 | 8.0 |

| 32-bit floating point | 1.1 |

| 32-bit floating point | 2.0 |

| 32-bit floating point float2 and float4 | 9.0 |

| 64-bit floating point | 6.0 |

| Data type | Supported since for shared memory |

|---|---|

| 32-bit integer signed/unsigned | 1.2 |

| 64-bit integer signed/unsigned | 2.0 |

| 32-bit floating point | 1.2 |

Note: Any missing lines or empty entries do reflect some lack of information on that exact item.cite-ref-61[61]

Tensor cores

| FMA per cycle per tensor core | Supported since | |

|---|---|---|

| Data Type | For dense matrices | For sparse matrices |

| 1-bit values (AND) | 8.0 as experimental | No |

| 1-bit values (XOR) | 7.5–8.9 as experimental | No |

| 4-bit integers | 7.5–8.9 as experimental | 8.0–8.9 as experimental |

| 4-bit floating point FP4 (E2M1) | 10.0 | |

| 6-bit floating point FP6 (E3M2 and E2M3) | 10.0 | |

| 8-bit integers | 7.2 | 8.0 |

| 8-bit floating point FP8 (E4M3 and E5M2… | 8.9 | |

| 16-bit floating point FP16 with FP16 ac… | 7.0 | 8.0 |

| 16-bit floating point FP16 with FP32 ac… | 7.0 | 8.0 |

| 16-bit floating point BF16 with FP32 ac… | 7.5 | 8.0 |

| 32-bit (19 bits used) floating point TF… | 7.5 | 8.0 |

| 64-bit floating point | 8.0 | No |

| FMA per cycle per tensor core | 7.0 | 7.2 | 7.5 Workstation |

|---|---|---|---|

| 1-bit values (AND) | No | | |

| 1-bit values (XOR) | No | | 1024 |

| 4-bit integers | No | | 256 |

| 4-bit floating point FP4 (E2M1) | No | | |

| 6-bit floating point FP6 (E3M2 and E2M3) | No | | |

| 8-bit integers | No | 128 | 128 |

| 8-bit floating point FP8 (E4M3 and E5M2… | No | | |

| 16-bit floating point FP16 with FP16 ac… | 64 | | 64 |

| 16-bit floating point FP16 with FP32 ac… | 64 | | 64 |

| 16-bit floating point BF16 with FP32 ac… | No | | 64 |

| 64-bit floating point | No | | No |

|---|---|---|---|

| 1-bit values (AND) | | 4096 | |

| 1-bit values (XOR) | | 4096 | |

| 4-bit integers | | 1024 | |

| 8-bit integers | | 512 | |

| 16-bit floating point FP16 with FP16 ac… | 64 | 256 | |

| 16-bit floating point FP16 with FP32 ac… | 32 | 256 | |

| 16-bit floating point BF16 with FP32 ac… | 32 | 256 | |

| 32-bit (19 bits used) floating point TF… | 32 | 128 | |

|---|---|---|---|

| 1-bit values (AND) | 2048 | | |

| 1-bit values (XOR) | 2048 | | |

| 4-bit integers | 512 | | |

| 8-bit integers | 256 | | |

| 8-bit floating point FP8 (E4M3 and E5M2… | | 256 | |

| 8-bit floating point FP8 (E4M3 and E5M2… | | 128 | |

| 16-bit floating point FP16 with FP16 ac… | 128 | | |

| 16-bit floating point FP16 with FP32 ac… | 64 | | 128 |

| 16-bit floating point BF16 with FP32 ac… | 64 | | 128 |

| 32-bit (19 bits used) floating point TF… | 32 | | 64 |

| FMA per cycle per tensor core | 9.0 | 10.0 | 10.1 |

|---|---|---|---|

| 4-bit floating point FP4 (E2M1) | | 4096 | tbd |

| 6-bit floating point FP6 (E3M2 and E2M3) | | 2048 | tbd |

| 8-bit integers | 1024 | 2048 | tbd |

| 8-bit floating point FP8 (E4M3 and E5M2… | 1024 | 2048 | tbd |

| 16-bit floating point FP16 with FP16 ac… | 512 | 1024 | tbd |

| 16-bit floating point FP16 with FP32 ac… | 512 | 1024 | tbd |

| 16-bit floating point BF16 with FP32 ac… | 512 | 1024 | tbd |

| 32-bit (19 bits used) floating point TF… | 256 | 512 | tbd |

| 64-bit floating point | 32 | 16 | tbd |

| FMA per cycle per tensor core | 12.0 |

|---|---|

| 4-bit floating point FP4 (E2M1) | 512 |

| 6-bit floating point FP6 (E3M2 and E2M3) | tbd |

| 8-bit integers | 256 |

| 8-bit floating point FP8 (E4M3 and E5M2… | 256 |

| 8-bit floating point FP8 (E4M3 and E5M2… | 128 |

| 16-bit floating point FP16 with FP16 ac… | 128 |

| 16-bit floating point FP16 with FP32 ac… | 64 |

| 16-bit floating point BF16 with FP32 ac… | 64 |

| 32-bit (19 bits used) floating point TF… | 32 |

| 64-bit floating point | tbd |

Note: Any missing lines or empty entries do reflect some lack of information on that exact item.cite-ref-65[65]cite-ref-66[66] cite-ref-67[67] cite-ref-68[68] cite-ref-69[69] cite-ref-70[70]

| Tensor Core Composition | 7.0 | 7.2, 7.5 | 8.0, 8.6 | 8.7 | 9.0 |

|---|---|---|---|---|---|

| Dot Product Unit Width in FP16 units (i… | 4 (8) | | 8 (16) | 4 (8) | 16 (32) |

| Dot Product Units per Tensor Core | 16 | | 32 | | |

| Tensor Cores per SM partition | 2 | | 1 | | |

| Full throughput (Bytes/cycle) per SM pa… | 256 | | 512 | 256 | 1024 |

| FP Tensor Cores: Minimum cycles for war… | 8 | | 4 | 8 | |

| FP Tensor Cores: Minimum Matrix Shape f… | 2048 | | | | |

| INT Tensor Cores: Minimum cycles for wa… | No | 4 | | | |

| INT Tensor Cores: Minimum Matrix Shape… | No | 1024 | 2048 | 1024 | |

cite-ref-78[78]cite-ref-79[79]cite-ref-80[80]cite-ref-81[81]

| FP64 Tensor Core Composition | 8.0 | 8.6 | 9.0 |

|---|---|---|---|

| Dot Product Unit Width in FP64 units (i… | 4 (32) | tbd | 4 (32) |

| Dot Product Units per Tensor Core | 4 | tbd | 8 |

| Tensor Cores per SM partition | 1 | | |

| Full throughput (Bytes/cycle) per SM pa… | 128 | tbd | 256 |

| Minimum cycles for warp-wide matrix cal… | 16 | tbd | |

| Minimum Matrix Shape for full throughpu… | 2048 | | |

Technical specifications

|---|---|---|---|

| Technical specifications | 1.0 | 1.1 | 1.2 |

| Maximum number of resident grids per de… | 1 | | |

| Maximum dimensionality of grid of threa… | 2 | | |

| Maximum x-dimension of a grid of thread… | 65535 | | |

| Maximum y-, or z-dimension of a grid of… | 65535 | | |

| Maximum dimensionality of thread block | 3 | | |

| Maximum x- or y-dimension of a block | 512 | | |

| Maximum z-dimension of a block | 64 | | |

| Maximum number of threads per block | 512 | | |

| Warp size | 32 | | |

| Maximum number of resident blocks per m… | 8 | | |

| Maximum number of resident warps per mu… | 24 | | 32 |

| Maximum number of resident threads per… | 768 | | 1024 |

| Number of 32-bit regular registers per… | 8 K | | 16 K |

| Number of 32-bit uniform registers per… | No | | |

| Maximum number of 32-bit registers per… | 8 K | | 16 K |

| Maximum number of 32-bit regular regist… | 124 | | |

| Maximum number of 32-bit uniform regist… | No | | |

| Amount of shared memory per multiproces… | 16 KiB | | |

| Maximum amount of shared memory per thr… | 16 KiB | | |

| Number of shared memory banks | 16 | | |

| Amount of local memory per thread | 16 KiB | | |

| Constant memory size accessible by CUDA… | 64 KiB | | |

| Cache working set per multiprocessor fo… | 8 KiB | | |

| Maximum width for 1D texture reference… | 8192 | | |

| Maximum width for 1D texture reference… | 2 27 | | |

| Maximum width and number of layers for… | 8192 × 512 | | |

| Maximum width and height for 2D texture… | 65536 × 32768 | | |

| Maximum width and height for 2D texture… | 65000 x 65000 | | |

| Maximum width and height for 2D texture… | — | | |

| Maximum width, height, and number of la… | 8192 × 8192 × 512 | | |

| Maximum width, height and depth for a 3… | 2048 3 | | |

| Maximum width (and height) for a cubema… | — | | |

| Maximum width (and height) and number o… | — | | |

| Maximum number of textures that can be… | 128 | | |

| Maximum number of Thread Blocks per Thr… | No | | |

| Technical specifications | 1.0 | 1.1 | 1.2 |

| Technical specifications | | |

|---|---|---|

| Technical specifications | 1.3 | 2.x |

| Maximum number of resident grids per de… | | 16 |

| Maximum dimensionality of grid of threa… | | 3 |

| Maximum x- or y-dimension of a block | | 1024 |

| Maximum number of threads per block | | 1024 |

| Maximum number of resident warps per mu… | | 48 |

| Maximum number of resident threads per… | | 1536 |

| Number of 32-bit regular registers per… | | 32 K |

| Maximum number of 32-bit registers per… | | 32 K |

| Maximum number of 32-bit regular regist… | | 63 |

| Amount of shared memory per multiproces… | | 16 / 48 KiB (of 64 KiB) |

| Maximum amount of shared memory per thr… | | 48 KiB |

| Number of shared memory banks | | 32 |

| Amount of local memory per thread | | 512 KiB |

| Cache working set per multiprocessor fo… | 24 KiB per TPC | 12 KiB |

| Maximum width for 1D texture reference… | | 65536 |

| Maximum width and number of layers for… | | 16384 × 2048 |

| Maximum width and height for 2D texture… | | 65536 × 65535 |

| Maximum width and height for 2D texture… | | 16384 x 16384 |

| Maximum width, height, and number of la… | | 16384 × 16384 × 2048 |

| Maximum width (and height) for a cubema… | | 16384 |

| Maximum width (and height) and number o… | | 16384 × 2046 |

| Maximum width for a 1D surface referenc… | | 65536 |

| Maximum width and number of layers for… | | 65536 × 2048 |

| Maximum width and height for a 2D surfa… | | 65536 × 32768 |

| Maximum width, height, and number of la… | | 65536 × 32768 × 2048 |

| Maximum width, height, and depth for a… | | 65536 × 32768 × 2048 |

| Maximum width (and height) for a cubema… | | 32768 |

| Maximum width and number of layers for… | | 32768 × 2046 |

| Maximum number of surfaces that can be… | | 8 |

| Maximum number of instructions per kern… | | 512 million |

| Technical specifications | 1.3 | 2.x |

|---|---|---|---|

| Technical specifications | 3.0 | 3.2 | 3.5 |

| Maximum number of resident grids per de… | | 4 | 32 |

| Maximum x-dimension of a grid of thread… | 2 31 − 1 | | |

| Maximum number of resident blocks per m… | 16 | | |

| Maximum number of resident warps per mu… | 64 | | |

| Maximum number of resident threads per… | 2048 | | |

| Number of 32-bit regular registers per… | 64 K | | |

| Maximum number of 32-bit registers per… | 64 K | 32 K | 64 K |

| Maximum number of 32-bit regular regist… | | 255 | |

| Cache working set per multiprocessor fo… | 12 – 48 KiB | | |

| Maximum width, height and depth for a 3… | 4096 3 | | |

| Maximum number of textures that can be… | 256 | | |

| Maximum number of surfaces that can be… | 16 | | |

| Technical specifications | 3.0 | 3.2 | 3.5 |

| Technical specifications | | |

|---|---|---|

| Technical specifications | 3.7 | 5.0 |

| Maximum number of resident blocks per m… | | 32 |

| Number of 32-bit regular registers per… | 128 K | 64 K |

| Amount of shared memory per multiproces… | 80 / 96 / 112 KiB (of 128 KiB) | 64 KiB |

| Cache working set per multiprocessor fo… | | 24 KiB |

| Maximum width and height for 2D texture… | | 65536 x 65536 |

| Maximum width for a 1D surface referenc… | | 16384 |

| Maximum width and number of layers for… | | 16384 × 2048 |

| Maximum width and height for a 2D surfa… | | 16384 × 65536 |

| Maximum width, height, and number of la… | | 16384 × 16384 × 2048 |

| Maximum width, height, and depth for a… | | 4096 × 4096 × 4096 |

| Maximum width (and height) for a cubema… | | 16384 |

| Maximum width and number of layers for… | | 16384 × 2046 |

| Technical specifications | 3.7 | 5.0 |

|---|---|---|---|---|

| Technical specifications | 5.2 | 5.3 | 6.0 | 6.1 |

| Maximum number of resident grids per de… | | 16 | 128 | 32 |

| Maximum number of 32-bit registers per… | | 32 K | 64 K | |

| Amount of shared memory per multiproces… | 96 KiB | 64 KiB | | 96 KiB |

| Cache working set per multiprocessor fo… | | | 4 KiB | 8 KiB |

| Cache working set per multiprocessor fo… | 48 KiB | 32 KiB | 24 KiB | 48 KiB |

| Maximum width for 1D texture reference… | | | 131072 | |

| Maximum width for 1D texture reference… | | | 2 28 | 2 27 |

| Maximum width and number of layers for… | | | 32768 x 2048 | |

| Maximum width and height for 2D texture… | | | 131072 x 65536 | |

| Maximum width and height for 2D texture… | | | 131072 x 65000 | |

| Maximum width and height for 2D texture… | | | 32768 x 32768 | |

| Maximum width, height, and number of la… | | | 32768 x 32768 x 2048 | |

| Maximum width, height and depth for a 3… | | | 16384 3 | |

| Maximum width (and height) for a cubema… | | | 32768 | |

| Maximum width (and height) and number o… | | | 32768 × 2046 | |

| Maximum width for a 1D surface referenc… | | | 32768 | |

| Maximum width and number of layers for… | | | 32768 × 2048 | |

| Maximum width and height for a 2D surfa… | | | 131072 × 65536 | |

| Maximum width, height, and number of la… | | | 32768 × 32768 × 2048 | |

| Maximum width, height, and depth for a… | | | 16384 × 16384 × 16384 | |

| Maximum width (and height) for a cubema… | | | 32768 | |

| Maximum width and number of layers for… | | | 32768 × 2046 | |

| Technical specifications | 5.2 | 5.3 | 6.0 | 6.1 |

| Technical specifications | | |

|---|---|---|

| Technical specifications | 6.2 | 7.0 |

| Maximum number of resident grids per de… | 16 | 128 |

| Maximum number of 32-bit registers per… | 32 K | 64 K |

| Amount of shared memory per multiproces… | 64 KiB | 0 / 8 / 16 / 32 / 64 / 96 KiB (of 128 K… |

| Maximum amount of shared memory per thr… | | 96 KiB |

| Cache working set per multiprocessor fo… | 24 KiB | 32 – 128 KiB |

| Maximum width for 1D texture reference… | | 2 28 |

| Maximum number of surfaces that can be… | | 32 |

| Technical specifications | 6.2 | 7.0 |

| Technical specifications | | |

|---|---|---|

| Technical specifications | 7.2 | 7.5 |

| Maximum number of resident grids per de… | 16 | 128 |

| Maximum number of resident blocks per m… | | 16 |

| Maximum number of resident warps per mu… | | 32 |

| Maximum number of resident threads per… | | 1024 |

| Number of 32-bit uniform registers per… | | 2 K |

| Maximum number of 32-bit uniform regist… | | 63 |

| Amount of shared memory per multiproces… | | 32 / 64 KiB (of 96 KiB) |

| Maximum amount of shared memory per thr… | 48 KiB | 64 KiB |

| Cache working set per multiprocessor fo… | | 32 – 64 KiB |

| Maximum width for 1D texture reference… | 2 27 | 2 28 |

| Technical specifications | 7.2 | 7.5 |

| Technical specifications | |

|---|---|

| Technical specifications | 8.0 |

| Maximum number of resident blocks per m… | 32 |

| Maximum number of resident warps per mu… | 64 |

| Maximum number of resident threads per… | 2048 |

| Amount of shared memory per multiproces… | 0 / 8 / 16 / 32 / 64 / 100 / 132 / 164… |

| Maximum amount of shared memory per thr… | 163 KiB |

| Cache working set per multiprocessor fo… | 28 – 192 KiB |

| Technical specifications | 8.0 |

| Technical specifications | |

|---|---|

| Technical specifications | 8.6 |

| Maximum number of resident blocks per m… | 16 |

| Maximum number of resident warps per mu… | 48 |

| Maximum number of resident threads per… | 1536 |

| Amount of shared memory per multiproces… | 0 / 8 / 16 / 32 / 64 / 100 KiB (of 128… |

| Maximum amount of shared memory per thr… | 99 KiB |

| Cache working set per multiprocessor fo… | 28 – 128 KiB |

| Technical specifications | 8.6 |

| Technical specifications | |

|---|---|

| Technical specifications | 8.7 |

| Amount of shared memory per multiproces… | 0 / 8 / 16 / 32 / 64 / 100 / 132 / 164… |

| Maximum amount of shared memory per thr… | 163 KiB |

| Cache working set per multiprocessor fo… | 28 – 192 KiB |

| Technical specifications | 8.7 |

| Technical specifications | |

|---|---|

| Technical specifications | 8.9 |

| Maximum number of resident blocks per m… | 24 |

| Amount of shared memory per multiproces… | 0 / 8 / 16 / 32 / 64 / 100 KiB (of 128… |

| Maximum amount of shared memory per thr… | 99 KiB |

| Cache working set per multiprocessor fo… | 28 – 128 KiB |

| Technical specifications | 8.9 |

| Technical specifications | | |

|---|---|---|

| Technical specifications | 9.0 | 10.x |

| Maximum number of resident blocks per m… | 32 | |

| Maximum number of resident warps per mu… | 64 | |

| Maximum number of resident threads per… | 2048 | |

| Amount of shared memory per multiproces… | 0 / 8 / 16 / 32 / 64 / 100 / 132 / 164… | |

| Maximum amount of shared memory per thr… | 227 KiB | |

| Cache working set per multiprocessor fo… | 28 – 256 KiB | |

| Maximum number of Thread Blocks per Thr… | 16 | |

| Technical specifications | 9.0 | 10.x |

| Technical specifications | |

|---|---|

| Technical specifications | 12.x |

| Maximum number of resident warps per mu… | 48 |

| Maximum number of resident threads per… | 1536 |

| Amount of shared memory per multiproces… | 0 / 8 / 16 / 32 / 64 / 100 KiB (of 128… |

| Maximum amount of shared memory per thr… | 99 KiB |

| Maximum number of Thread Blocks per Thr… | 8 |

| Technical specifications | 12.x |

Multiprocessor architecture

| Architecture specifications | Compute capability (version) | |

|---|---|---|

| Architecture specifications | 1.0 | 1.1 |

| Number of ALU lanes for INT32 arithmeti… | 8 | |

| Number of ALU lanes for any INT32 or FP… | 8 | |

| Number of ALU lanes for FP32 arithmetic… | 8 | |

| Number of ALU lanes for FP16x2 arithmet… | No | |

| Number of ALU lanes for FP64 arithmetic… | No | |

| Number of Load/Store Units | 4 per 2 SM | 8 per 2 SM |

| Number of special function units for si… | 2 | |

| Number of texture mapping units (TMU) | 4 per 2 SM | 8 per 2 SM |

| Number of ALU lanes for uniform INT32 a… | No | |

| Number of tensor cores | No | |

| Number of raytracing cores | No | |

| Number of SM Partitions = Processing Bl… | 1 | |

| Number of warp schedulers per SM partit… | 1 | |

| Max number of new instructions issued e… | 2 | |

| Size of unified memory for data cache a… | 16 KiB | |

| Size of L3 instruction cache per GPU | | |

| Size of L2 instruction cache per Textur… | | |

| Size of L1.5 instruction cache per SM | | |

| Size of L1 instruction cache per SM | | |

| Size of L0 instruction cache per SM par… | only 1 partition per SM | |

| Instruction Width | 32 bits instructions and 64 bits instru… | |

| Memory Bus Width per Memory Partition i… | 64 ((G)DDR) | |

| L2 Cache per Memory Partition | 16 KiB | |

| Number of Render Output Units (ROP) per… | 4 | |

| Architecture specifications | 1.0 | 1.1 |

| Architecture specifications | Compute capability (version) | |

|---|---|---|---|

| Architecture specifications | 1.2 | 1.3 | 2.0 |

| Number of ALU lanes for INT32 arithmeti… | | | 32 |

| Number of ALU lanes for any INT32 or FP… | | | 32 |

| Number of ALU lanes for FP32 arithmetic… | | | 32 |

| Number of special function units for si… | | | 4 |

| Number of warp schedulers per SM partit… | | | 2 |

| Max number of new instructions issued e… | | | 1 |

| Size of unified memory for data cache a… | | 16 KiB | 64 KiB |

| Size of L3 instruction cache per GPU | | 32 KiB | |

| Size of L2 instruction cache per Textur… | | 8 KiB | |

| Size of L1.5 instruction cache per SM | | 4 KiB | |

| Size of L1 instruction cache per SM | | 4 KiB | |

| L2 Cache per Memory Partition | | 32 KiB | 128 KiB |

| Number of Render Output Units (ROP) per… | | | 8 |

| Architecture specifications | 1.2 | 1.3 | 2.0 |

| Architecture specifications | | |

|---|---|---|

| Architecture specifications | 2.1 | 3.0 |

| Number of ALU lanes for INT32 arithmeti… | 48 | 192 |

| Number of ALU lanes for any INT32 or FP… | 48 | 192 |

| Number of ALU lanes for FP32 arithmetic… | 48 | 192 |

| Number of ALU lanes for FP64 arithmetic… | 4 by FP32 | 8 |

| Number of Load/Store Units | | 32 |

| Number of special function units for si… | 8 | 32 |

| Number of texture mapping units (TMU) | 4 / 8 | 16 |

| Number of warp schedulers per SM partit… | | 4 |

| Max number of new instructions issued e… | 2 | 2 |

| Size of L3 instruction cache per GPU | | use L2 Data Cache |

| Size of L2 instruction cache per Textur… | | use L2 Data Cache |

| Instruction Width | | 64 bits instructions + 64 bits control… |

| Architecture specifications | 2.1 | 3.0 |

|---|---|---|---|

| Architecture specifications | 3.2 | 3.5 | 3.7 |

| Number of ALU lanes for FP64 arithmetic… | | 8 / 64 | 64 |

| Number of texture mapping units (TMU) | 8 | 16 | |

| Size of unified memory for data cache a… | | | 128 KiB |

| Size of L1.5 instruction cache per SM | | | 32 KiB |

| Size of L1 instruction cache per SM | | | 8 KiB |

| L2 Cache per Memory Partition | | 256 KiB | |

| Number of Render Output Units (ROP) per… | 4 | 8 | |

| Architecture specifications | 3.2 | 3.5 | 3.7 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 5.0 |

| Number of ALU lanes for INT32 arithmeti… | 128 |

| Number of ALU lanes for any INT32 or FP… | 128 |

| Number of ALU lanes for FP32 arithmetic… | 128 |

| Number of ALU lanes for FP64 arithmetic… | 4 |

| Number of texture mapping units (TMU) | 8 |

| Number of SM Partitions = Processing Bl… | 4 |

| Number of warp schedulers per SM partit… | 1 |

| Size of unified memory for data cache a… | 64 KiB SM + 24 KiB L1 (separate) |

| Size of L0 instruction cache per SM par… | No |

| Instruction Width | 64 bits instructions + 64 bits control… |

| L2 Cache per Memory Partition | 1 MiB |

| Architecture specifications | 5.0 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 5.2 |

| Size of unified memory for data cache a… | 96 KiB SM + 24 KiB L1 (separate) |

| Size of L1.5 instruction cache per SM | 32 KiB |

| L2 Cache per Memory Partition | 512 KiB |

| Number of Render Output Units (ROP) per… | 16 |

| Architecture specifications | 5.2 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 5.3 |

| Number of ALU lanes for INT32 arithmeti… | 128 |

| Number of ALU lanes for any INT32 or FP… | 128 |

| Number of ALU lanes for FP32 arithmetic… | 128 |

| Number of ALU lanes for FP16x2 arithmet… | 128 |

| Number of special function units for si… | 16 |

| Size of unified memory for data cache a… | 64 KiB SM + 24 KiB L1 (separate) |

| Size of L1.5 instruction cache per SM | 48 KiB |

| Memory Bus Width per Memory Partition i… | 32 ((G)DDR) |

| L2 Cache per Memory Partition | 128 KiB |

| Number of Render Output Units (ROP) per… | 8 |

| Architecture specifications | 5.3 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 6.0 |

| Number of ALU lanes for INT32 arithmeti… | 64 |

| Number of ALU lanes for any INT32 or FP… | 64 |

| Number of ALU lanes for FP32 arithmetic… | 64 |

| Number of ALU lanes for FP16x2 arithmet… | 64 |

| Number of ALU lanes for FP64 arithmetic… | 32 |

| Number of Load/Store Units | 16 |

| Number of SM Partitions = Processing Bl… | 2 |

| Size of unified memory for data cache a… | 64 KiB SM + 24 KiB L1 (separate) |

| Size of L1.5 instruction cache per SM | 128 KiB |

| Size of L1 instruction cache per SM | 8 KiB |

| Memory Bus Width per Memory Partition i… | 512 (HBM) |

| L2 Cache per Memory Partition | 512 KiB |

| Number of Render Output Units (ROP) per… | 12 |

| Architecture specifications | 6.0 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 6.1 |

| Number of ALU lanes for INT32 arithmeti… | 128 |

| Number of ALU lanes for any INT32 or FP… | 128 |

| Number of ALU lanes for FP32 arithmetic… | 128 |

| Number of ALU lanes for FP16x2 arithmet… | 1 |

| Number of ALU lanes for FP64 arithmetic… | 4 |

| Number of Load/Store Units | 32 |

| Number of special function units for si… | 32 |

| Number of SM Partitions = Processing Bl… | 4 |

| Size of unified memory for data cache a… | 96 KiB SM + 24 KiB L1 (separate) |

| Size of L1.5 instruction cache per SM | 32 KiB |

| Memory Bus Width per Memory Partition i… | 32 ((G)DDR) |

| L2 Cache per Memory Partition | 256 KiB |

| Number of Render Output Units (ROP) per… | 8 |

| Architecture specifications | 6.1 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 6.2 |

| Number of ALU lanes for INT32 arithmeti… | 128 |

| Number of ALU lanes for any INT32 or FP… | 128 |

| Number of ALU lanes for FP32 arithmetic… | 128 |

| Number of ALU lanes for FP16x2 arithmet… | 128 |

| Size of unified memory for data cache a… | 64 KiB SM + 24 KiB L1 (separate) |

| L2 Cache per Memory Partition | 128 KiB |

| Number of Render Output Units (ROP) per… | 4 |

| Architecture specifications | 6.2 |

| Architecture specifications | |

|---|---|

| Architecture specifications | 7.0 |

| Number of ALU lanes for INT32 arithmeti… | 64 |

| Number of ALU lanes for any INT32 or FP… | — |

| Number of ALU lanes for FP32 arithmetic… | 64 |

| Number of ALU lanes for FP16x2 arithmet… | 64 |

| Number of ALU lanes for FP64 arithmetic… | 32 |

| Number of special function units for si… | 16 |

| Number of texture mapping units (TMU) | 4 |

| Number of tensor cores | 8 (1st gen.) |

| Max number of new instructions issued e… | 1 |

| Size of unified memory for data cache a… | 128 KiB |

| Size of L1.5 instruction cache per SM | 128 KiB |

| Size of L1 instruction cache per SM | 128 KiB |

| Size of L0 instruction cache per SM par… | 12 KiB |

| Instruction Width | 128 bits combined instruction and contr… |

| Memory Bus Width per Memory Partition i… | 512 (HBM) |

| L2 Cache per Memory Partition | 768 KiB |

| Number of Render Output Units (ROP) per… | 16 |

| Architecture specifications | 7.0 |

|---|---|---|---|

| Architecture specifications | 7.2 | 7.5 | 8.0 |

| Number of ALU lanes for FP32 arithmetic… | | | 64 |

| Number of ALU lanes for FP16x2 arithmet… | | | 128 |

| Number of ALU lanes for FP64 arithmetic… | | 2 | 32 |

| Number of Load/Store Units | | | 16 |

| Number of ALU lanes for uniform INT32 a… | | 2 | |

| Size of unified memory for data cache a… | | 96 KiB | 192 KiB |

| Size of L1.5 instruction cache per SM | | ~46 KiB | 128 KiB |

| Size of L1 instruction cache per SM | | ~46 KiB | 128 KiB |

| Size of L0 instruction cache per SM par… | | 16 KiB? | 32 KiB |

| L2 Cache per Memory Partition | 64 KiB | 512 KiB | 4 MiB |

| Number of Render Output Units (ROP) per… | 2 | 8 | 16 |

| Architecture specifications | 7.2 | 7.5 | 8.0 |

|---|---|---|---|---|

| Architecture specifications | 8.6 | 8.7 | 8.9 | 9.0 |

| Number of ALU lanes for INT32 arithmeti… | 64 | | 64 | |

| Number of ALU lanes for any INT32 or FP… | 64 | | — | |

| Number of ALU lanes for FP32 arithmetic… | | | 128 | 128 |

| Number of ALU lanes for FP16x2 arithmet… | 128 | | 64 | 128 |

| Number of ALU lanes for FP64 arithmetic… | 2 | | | 64 |

| Number of Load/Store Units | | | | 32 |

| Size of unified memory for data cache a… | 128 KiB | 192 KiB | 128 KiB | 256 KiB |

| L2 Cache per Memory Partition | 512 KiB | | 8 MiB | 5 MiB |

| Architecture specifications | 8.6 | 8.7 | 8.9 | 9.0 |

| Architecture specifications | | |

|---|---|---|

| Architecture specifications | 10.x | 12.x |

| Number of ALU lanes for INT32 arithmeti… | 128 | |

| Number of ALU lanes for any INT32 or FP… | 128 | |

| Number of ALU lanes for FP32 arithmetic… | 128 | |

| Number of ALU lanes for FP16x2 arithmet… | 128 | |

| Number of ALU lanes for FP64 arithmetic… | | 2 |

| Memory Bus Width per Memory Partition i… | | 32 ((G)DDR) |

| L2 Cache per Memory Partition | 6.25 MiB | 8 MiB |

| Architecture specifications | 10.x | 12.x |

For more information read the Nvidia CUDA C++ Programming Guide.cite-ref-117[117]

Usages of CUDA architecture

• Accelerated rendering of 3D graphics

• Accelerated interconversion of video file formats

• Accelerated encryption, decryption and compression

• Bioinformatics, e.g. NGS DNA sequencing BarraCUDAcite-ref-118[118]

• Distributed calculations, such as predicting the native conformation of proteins

• Medical analysis simulations, for example virtual reality based on CT and MRI scan images

• Physical simulations,cite-ref-119[119] particularly in fluid dynamics

• Neural network training in machine learning problems

• Large Language Model inference

• Face recognition

• Volunteer computing projects, such as SETI@home and other projects using BOINC software

• Molecular dynamics

• Mining cryptocurrencies

• Structure from motion (SfM) software

Comparison with competitors

CUDA competes with other GPU computing stacks: Intel OneAPI and AMD ROCm.

Whereas Nvidia's CUDA is closed-source, Intel's OneAPI and AMD's ROCm are open source.

Intel OneAPI

oneAPI is an initiative based in open standards, created to support software development for multiple hardware architectures.cite-ref-120[120] The oneAPI libraries must implement open specifications that are discussed publicly by the Special Interest Groups, offering the possibility for any developer or organization to implement their own versions of oneAPI libraries.cite-ref-121[121]cite-ref-122[122]

Originally made by Intel, other hardware adopters include Fujitsu and Huawei.

Unified Acceleration Foundation (UXL)

Unified Acceleration Foundation (UXL) is a new technology consortium working on the continuation of the OneAPI initiative, with the goal to create a new open standard accelerator software ecosystem, related open standards and specification projects through Working Groups and Special Interest Groups (SIGs). The goal is to offer open alternatives to Nvidia's CUDA. The main companies behind it are Intel, Google, ARM, Qualcomm, Samsung, Imagination, and VMware.cite-ref-123[123]

AMD ROCm

ROCmcite-ref-124[124] is an open source software stack for graphics processing unit (GPU) programming from Advanced Micro Devices (AMD).

See also

• SYCL – an open standard from Khronos Group for programming a variety of platforms, including GPUs, with single-source modern C++, similar to higher-level CUDA Runtime API (single-source)

• BrookGPU – the Stanford University graphics group's compiler

• Array programming

• Parallel computing

• Stream processing

• rCUDA – an API for computing on remote computers

• Molecular modeling on GPUs

• Vulkan – low-level, high-performance 3D graphics and computing API

• OptiX – ray tracing API by NVIDIA

• CUDA binary (cubin) – a type of fat binary

• Numerical Library Collection – by NEC for their vector processor

References

cite-note-11. "NVIDIA® CUDA™ Unleashes Power of GPU Computing - Press Release". nvidia.com. Archived from the original on 29 March 2007. Retrieved 26 January 2025.

cite-note-0-22. ↑ citerefshahShah, Agam. "Nvidia not totally against third parties making CUDA chips". www.theregister.com. Retrieved 2024-04-25.

cite-note-33. ↑ "Nvidia CUDA Home Page". 18 July 2017.

cite-note-cuda-intro-anandtech-44. ↑ citerefshimpiwilson2006Shimpi, Anand Lal; Wilson, Derek (November 8, 2006). "Nvidia's GeForce 8800 (G80): GPUs Re-architected for DirectX 10". AnandTech. Archived from the original on April 24, 2010. Retrieved May 16, 2015.

cite-note-55. ↑ "Introduction — nsight-visual-studio-edition 12.6 documentation". docs.nvidia.com. Retrieved 2024-10-10.

cite-note-cuda-intro-tomshardware-66. ↑ citerefabi-chahla2008Abi-Chahla, Fedy (June 18, 2008). "Nvidia's CUDA: The End of the CPU?". Tom's Hardware. Retrieved May 17, 2015.

cite-note-77. ↑ citerefjones2025Jones, Stephen (2025-04-22). What is CUDA? (Video). Computerphile. Retrieved 2025-07-24 – via YouTube.

cite-note-88. ↑ citerefzunitch2018Zunitch, Peter (2018-01-24). "CUDA vs. OpenCL vs. OpenGL". Videomaker. Retrieved 2018-09-16.

cite-note-99. ↑ "OpenCL". NVIDIA Developer. 2013-04-24. Retrieved 2019-11-04.

cite-note-1-1010. ↑ citerefcosgroveCosgrove, Emma. "Ian Buck built Nvidia's secret weapon. He may spend the rest of his career defending it". Business Insider. Retrieved 2025-07-24.

cite-note-1111. ↑ citerefwitt2023Witt, Stephen (2023-11-27). "How Jensen Huang's Nvidia Is Powering the A.I. Revolution". The New Yorker. ISSN 0028-792X. Retrieved 2023-12-10.

cite-note-1212. ↑ "CUDA LLVM Compiler". 7 May 2012.

cite-note-1313. ↑ First OpenCL demo on a GPU on YouTube

cite-note-1414. ↑ DirectCompute Ocean Demo Running on Nvidia CUDA-enabled GPU on YouTube

cite-note-ioannidis08-1515. ↑ citerefvasiliadisantonatospolychronakismarkatos2008Vasiliadis, Giorgos; Antonatos, Spiros; Polychronakis, Michalis; Markatos, Evangelos P.; Ioannidis, Sotiris (September 2008). "Gnort: High Performance Network Intrusion Detection Using Graphics Processors" (PDF). Recent Advances in Intrusion Detection. Lecture Notes in Computer Science. Vol. 5230. pp. 116–134. doi:10.1007/978-3-540-87403-4_7. ISBN 978-3-540-87402-7.

cite-note-1616. ↑ citerefschatztrapnelldelchervarshney2007Schatz, Michael C.; Trapnell, Cole; Delcher, Arthur L.; Varshney, Amitabh (2007). "High-throughput sequence alignment using Graphics Processing Units". BMC Bioinformatics. 8: 474. doi:10.1186/1471-2105-8-474. PMC 2222658. PMID 18070356.

cite-note-manavski2008-1717. ↑ citerefmanavskigiorgio2008Manavski, Svetlin A.; Giorgio, Valle (2008). "CUDA compatible GPU cards as efficient hardware accelerators for Smith-Waterman sequence alignment". BMC Bioinformatics. 10 (Suppl 2): S10. doi:10.1186/1471-2105-9-S2-S10. PMC 2323659. PMID 18387198.

cite-note-1818. ↑ "Pyrit – Google Code".

cite-note-1919. ↑ "Use your Nvidia GPU for scientific computing". BOINC. 2008-12-18. Archived from the original on 2008-12-28. Retrieved 2017-08-08.

cite-note-2020. ↑ "Nvidia CUDA Software Development Kit (CUDA SDK) – Release Notes Version 2.0 for MAC OS X". Archived from the original on 2009-01-06.

cite-note-2121. ↑ "CUDA 1.1 – Now on Mac OS X". February 14, 2008. Archived from the original on November 22, 2008.

cite-note-2222. ↑ "CUDA 11 Features Revealed". 14 May 2020.

cite-note-2323. ↑ "CUDA Toolkit 11.1 Introduces Support for GeForce RTX 30 Series and Quadro RTX Series GPUs". 23 September 2020.

cite-note-2424. ↑ "Enhancing Memory Allocation with New NVIDIA CUDA 11.2 Features". 16 December 2020.

cite-note-2525. ↑ "Exploring the New Features of CUDA 11.3". 16 April 2021.

cite-note-2626. ↑ citerefsilbersteinschustergeigerpatney2008Silberstein, Mark; Schuster, Assaf; Geiger, Dan; Patney, Anjul; Owens, John D. (2008). "Efficient computation of sum-products on GPUs through software-managed cache" (PDF). Proceedings of the 22nd annual international conference on Supercomputing – ICS '08 (PDF). Proceedings of the 22nd annual international conference on Supercomputing – ICS '08. pp. 309–318. doi:10.1145/1375527.1375572. ISBN 978-1-60558-158-3.

cite-note-cuda-prog-v8-2727. ↑ "CUDA C Programming Guide v8.0" (PDF). nVidia Developer Zone. January 2017. p. 19. Retrieved 22 March 2017.

cite-note-2828. ↑ "NVCC forces c++ compilation of .cu files". 29 November 2011.

cite-note-2929. ↑ citerefwhiteheadfit-floreaWhitehead, Nathan; Fit-Florea, Alex. "Precision & Performance: Floating Point and IEEE 754 Compliance for Nvidia GPUs" (PDF). Nvidia. Retrieved November 18, 2014.

cite-note-cuda-products-3030. ↑ "CUDA-Enabled Products". CUDA Zone. Nvidia Corporation. Retrieved 2008-11-03.

cite-note-3131. ↑ "Coriander Project: Compile CUDA Codes To OpenCL, Run Everywhere". Phoronix.

cite-note-3232. ↑ citerefperkins2017Perkins, Hugh (2017). "cuda-on-cl" (PDF). IWOCL. Retrieved August 8, 2017.

cite-note-3333. ↑ "hughperkins/coriander: Build NVIDIA® CUDA™ code for OpenCL™ 1.2 devices". GitHub. May 6, 2019.

cite-note-3434. ↑ "CU2CL Documentation". chrec.cs.vt.edu.

cite-note-3535. ↑ "GitHub – vosen/ZLUDA". GitHub.

cite-note-3636. ↑ citereflarabel2024Larabel, Michael (2024-02-12), "AMD Quietly Funded A Drop-In CUDA Implementation Built On ROCm: It's Now Open-Source", Phoronix, retrieved 2024-02-12

cite-note-3737. ↑ "GitHub – chip-spv/chipStar". GitHub.

cite-note-3838. ↑ "PyCUDA".

cite-note-3939. ↑ "pycublas". Archived from the original on 2009-04-20. Retrieved 2017-08-08.

cite-note-4040. ↑ "CuPy". Retrieved 2020-01-08.

cite-note-412. "NVIDIA CUDA Programming Guide. Version 1.0" (PDF). June 23, 2007.

cite-note-423. "NVIDIA CUDA Programming Guide. Version 2.1" (PDF). December 8, 2008.

cite-note-434. "NVIDIA CUDA Programming Guide. Version 2.2" (PDF). April 2, 2009.

cite-note-445. "NVIDIA CUDA Programming Guide. Version 2.2.1" (PDF). May 26, 2009.

cite-note-456. "NVIDIA CUDA Programming Guide. Version 2.3.1" (PDF). August 26, 2009.

cite-note-467. "NVIDIA CUDA Programming Guide. Version 3.0" (PDF). February 20, 2010.

cite-note-478. "NVIDIA CUDA C Programming Guide. Version 3.1.1" (PDF). July 21, 2010.

cite-note-489. "NVIDIA CUDA C Programming Guide. Version 3.2" (PDF). November 9, 2010.

cite-note-4910. "CUDA 11.0 Release Notes". NVIDIA Developer.

cite-note-5011. "CUDA 11.1 Release Notes". NVIDIA Developer.

cite-note-5112. "CUDA 11.5 Release Notes". NVIDIA Developer.

cite-note-5213. "CUDA 11.8 Release Notes". NVIDIA Developer.

cite-note-5314. "NVIDIA Quadro NVS 420 Specs". TechPowerUp GPU Database. 25 August 2023.

cite-note-5415. citereflarabel2017Larabel, Michael (March 29, 2017). "NVIDIA Rolls Out Tegra X2 GPU Support In Nouveau". Phoronix. Retrieved August 8, 2017.

cite-note-5516. Nvidia Xavier Specs on TechPowerUp (preliminary)

cite-note-5617. "Welcome — Jetson LinuxDeveloper Guide 34.1 documentation".

cite-note-5718. "NVIDIA Bringing up Open-Source Volta GPU Support for Their Xavier SoC".

cite-note-5819. "NVIDIA Ada Lovelace Architecture".

cite-note-5920. Dissecting the Turing GPU Architecture through Microbenchmarking

cite-note-6060. ↑ "H.1. Features and Technical Specifications – Table 13. Feature Support per Compute Capability". docs.nvidia.com. Retrieved 2020-09-23.

cite-note-6161. ↑ "CUDA C++ Programming Guide".

cite-note-6221. Fused-Multiply-Add, actually executed, Dense Matrix

cite-note-6322. as SASS since 7.5, as PTX since 8.0

cite-note-unofficial-support-in-sass-6423. unofficial support in SASS

cite-note-6565. ↑ "Technical brief. NVIDIA Jetson AGX Orin Series" (PDF). nvidia.com. Retrieved 5 September 2023.

cite-note-6666. ↑ "NVIDIA Ampere GA102 GPU Architecture" (PDF). nvidia.com. Retrieved 5 September 2023.

cite-note-6767. ↑ citerefluofanlidu2024Luo, Weile; Fan, Ruibo; Li, Zeyu; Du, Dayou; Wang, Qiang; Chu, Xiaowen (2024). "Benchmarking and Dissecting the Nvidia Hopper GPU Architecture". arXiv:2402.13499v1 [cs.AR].

cite-note-6868. ↑ "Datasheet NVIDIA A40" (PDF). nvidia.com. Retrieved 27 April 2024.

cite-note-6969. ↑ "NVIDIA AMPERE GA102 GPU ARCHITECTURE" (PDF). 27 April 2024.

cite-note-7070. ↑ "Datasheet NVIDIA L40" (PDF). 27 April 2024.

cite-note-7124. In the Whitepapers the Tensor Core cube diagrams represent the Dot Product Unit Width into the height (4 FP16 for Volta and Turing, 8 FP16 for A100, 4 FP16 for GA102, 16 FP16 for GH100). The other two dimensions represent the number of Dot Product Units (4x4 = 16 for Volta and Turing, 8x4 = 32 for Ampere and Hopper). The resulting gray blocks are the FP16 FMA operations per cycle. Pascal without Tensor core is only shown for speed comparison as is Volta V100 with non-FP16 datatypes.

cite-note-7225. "NVIDIA Turing Architecture Whitepaper" (PDF). nvidia.com. Retrieved 5 September 2023.

cite-note-7326. "NVIDIA Tensor Core GPU" (PDF). nvidia.com. Retrieved 5 September 2023.

cite-note-7427. "NVIDIA Hopper Architecture In-Depth". 22 March 2022.

cite-note-referencec-7528. shape x converted operand size, e.g. 2 tensor cores x 4x4x4xFP16/cycle = 256 Bytes/cycle

cite-note-product-first-3-table-rows-7629. = product first 3 table rows

cite-note-referenced-7730. = product of previous 2 table rows; shape: e.g. 8x8x4xFP16 = 512 Bytes

cite-note-7878. ↑ citerefsunligengstuijk2023Sun, Wei; Li, Ang; Geng, Tong; Stuijk, Sander; Corporaal, Henk (2023). "Dissecting Tensor Cores via Microbenchmarks: Latency, Throughput and Numeric Behaviors". IEEE Transactions on Parallel and Distributed Systems. 34 (1): 246–261. arXiv:2206.02874. doi:10.1109/tpds.2022.3217824. S2CID 249431357.

cite-note-7979. ↑ "Parallel Thread Execution ISA Version 7.7".

cite-note-8080. ↑ citerefraihangoliaamodt2018Raihan, Md Aamir; Goli, Negar; Aamodt, Tor (2018). "Modeling Deep Learning Accelerator Enabled GPUs". arXiv:1811.08309 [cs.MS].

cite-note-8181. ↑ "NVIDIA Ada Lovelace Architecture".

cite-note-referencee-8231. citerefjiamaggionismithdaniele-paolo-scarpazza2019Jia, Zhe; Maggioni, Marco; Smith, Jeffrey; Daniele Paolo Scarpazza (2019). "Dissecting the NVidia Turing T4 GPU via Microbenchmarking". arXiv:1903.07486 [cs.DC].

cite-note-8332. citerefburgess2019Burgess, John (2019). "RTX ON – The NVIDIA TURING GPU". 2019 IEEE Hot Chips 31 Symposium (HCS). pp. 1–27. doi:10.1109/HOTCHIPS.2019.8875651. ISBN 978-1-7281-2089-8. S2CID 204822166.

cite-note-8433. citerefburgess2019Burgess, John (2019). "RTX ON – The NVIDIA TURING GPU". 2019 IEEE Hot Chips 31 Symposium (HCS). pp. 1–27. doi:10.1109/HOTCHIPS.2019.8875651. ISBN 978-1-7281-2089-8. S2CID 204822166.

cite-note-8534. dependent on device

cite-note-tegra-x1-8635. "Tegra X1". 9 January 2015.

cite-note-8736. NVIDIA H100 Tensor Core GPU Architecture

cite-note-8837. H.1. Features and Technical Specifications – Table 14. Technical Specifications per Compute Capability

cite-note-8938. NVIDIA Hopper Architecture In-Depth

cite-note-9039. can only execute 160 integer instructions according to programming guide

cite-note-9140. 128 according to [1]. 64 from FP32 + 64 separate units?

cite-note-9241. 64 by FP32 cores and 64 by flexible FP32/INT cores.

cite-note-9342. "CUDA C++ Programming Guide".

cite-note-9443. 32 FP32 lanes combine to 16 FP64 lanes. Maybe lower depending on model.

cite-note-9544. only supported by 16 FP32 lanes, they combine to 4 FP64 lanes

cite-note-depending-on-model-9645. depending on model

cite-note-9746. Effective speed, probably over FP32 ports. No description of actual FP64 cores.

cite-note-9847. Can also be used for integer additions and comparisons

cite-note-9948. 2 clock cycles/instruction for each SM partition citerefburgess2019Burgess, John (2019). "RTX ON – The NVIDIA TURING GPU". 2019 IEEE Hot Chips 31 Symposium (HCS). pp. 1–27. doi:10.1109/HOTCHIPS.2019.8875651. ISBN 978-1-7281-2089-8. S2CID 204822166.

cite-note-inside-volta-10049. citerefdurantgirouxharrisstam2017Durant, Luke; Giroux, Olivier; Harris, Mark; Stam, Nick (May 10, 2017). "Inside Volta: The World's Most Advanced Data Center GPU". Nvidia developer blog.

cite-note-10150. The schedulers and dispatchers have dedicated execution units unlike with Fermi and Kepler.

cite-note-10251. Dispatching can overlap concurrently, if it takes more than one cycle (when there are less execution units than 32/SM Partition)

cite-note-10352. Can dual issue MAD pipe and SFU pipe

cite-note-10453. No more than one scheduler can issue 2 instructions at once. The first scheduler is in charge of warps with odd IDs. The second scheduler is in charge of warps with even IDs.

cite-note-shared-memory-only-no-data-cache-10554. shared memory only, no data cache

cite-note-referencea-10655. shared memory separate, but L1 includes texture cache

cite-note-10756. "H.6.1. Architecture". docs.nvidia.com. Retrieved 2019-05-13.

cite-note-10857. "Demystifying GPU Microarchitecture through Microbenchmarking" (PDF).

cite-note-referencef-10958. citerefjiamaggionistaigerscarpazza2018Jia, Zhe; Maggioni, Marco; Staiger, Benjamin; Scarpazza, Daniele P. (2018). "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking". arXiv:1804.06826 [cs.DC].

cite-note-11059. citerefjiamaggionismithdaniele-paolo-scarpazza2019Jia, Zhe; Maggioni, Marco; Smith, Jeffrey; Daniele Paolo Scarpazza (2019). "Dissecting the NVidia Turing T4 GPU via Microbenchmarking". arXiv:1903.07486 [cs.DC].

cite-note-11160. "Dissecting the Ampere GPU Architecture through Microbenchmarking".

cite-note-11261. Note that citerefjiamaggionismithdaniele-paolo-scarpazza2019Jia, Zhe; Maggioni, Marco; Smith, Jeffrey; Daniele Paolo Scarpazza (2019). "Dissecting the NVidia Turing T4 GPU via Microbenchmarking". arXiv:1903.07486 [cs.DC]. disagrees and states 2 KiB L0 instruction cache per SM partition and 16 KiB L1 instruction cache per SM

cite-note-11362. "asfermi Opcode". GitHub.

cite-note-referenceb-11463. for access with texture engine only

cite-note-11564. 25% disabled on RTX 4060, RTX 4070, RTX 4070 Ti and RTX 4090

cite-note-11665. 25% disabled on RTX 5070 Ti and RTX 5090

cite-note-117117. ↑ "CUDA C++ Programming Guide, Compute Capabilities". docs.nvidia.com. Retrieved 2025-02-06.

cite-note-118118. ↑ "nVidia CUDA Bioinformatics: BarraCUDA". BioCentric. 2019-07-19. Retrieved 2019-10-15.

cite-note-119119. ↑ "Part V: Physics Simulation". NVIDIA Developer. Retrieved 2020-09-11.

cite-note-120120. ↑ "oneAPI Programming Model". oneAPI.io. Retrieved 2024-07-27.

cite-note-121121. ↑ "Specifications | oneAPI". oneAPI.io. Retrieved 2024-07-27.

cite-note-122122. ↑ "oneAPI Specification — oneAPI Specification 1.3-rev-1 documentation". oneapi-spec.uxlfoundation.org. Retrieved 2024-07-27.

cite-note-123123. ↑ "Exclusive: Behind the plot to break Nvidia's grip on AI by targeting software". Reuters. Retrieved 2024-04-05.

cite-note-124124. ↑ "Question: What does ROCm stand for? · Issue #1628 · RadeonOpenCompute/ROCm". Github.com. Retrieved January 18, 2022.